#include <iostream>
#include <cuda_runtime.h>
#include "../common/common.h"
#define DEVICENUM 0
#define BLOCKSIZE 64 // 配置一个threads=64的一维线程块，也就是只有两个线程束

__global__ void threadYesDiv(float *res)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float ia, ib;
    ia = ib = 0.0f;

    if (tid % 2 == 0)
    {
        ia = 100.0f;
    }
    else
    {
        ib = 200.0f;
    }

    res[tid] = ia + ib;
}

__global__ void threadNoDiv(float *res)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float ia, ib;
    ia = ib = 0.0f;

    // 注意：warpSize和threadIdx.x一样，可以直接被核函数使用，默认是32
    // printf("thread wrap size = %d\n", warpSize);
    if ((tid / warpSize) % 2 == 0)
    {
        ia = 100.0f;
    }
    else
    {
        ib = 200.0f;
    }

    res[tid] = ia + ib;
}

int main(int argc, char **argv)
{
    cudaSetDevice(DEVICENUM);

    int nElem = BLOCKSIZE;
    size_t nBytes = nElem * sizeof(float);

    float *device_res_array = nullptr;
    CHECK(cudaMalloc((float **)&device_res_array, nBytes));

    dim3 gridsize(1);
    dim3 blocksize(BLOCKSIZE);

    threadYesDiv<<<gridsize, blocksize>>>(device_res_array);
    CHECK(cudaDeviceSynchronize());

    threadNoDiv<<<gridsize, blocksize>>>(device_res_array);
    CHECK(cudaDeviceSynchronize());

    cudaFree(device_res_array);
    return 0;
}